/**
 * @file test_kernel.cpp
 * @author 逆流 (1171267147@qq.com)
 * @brief 测试OpenCL内核创建、管理和执行功能
 * @version 0.2
 * @date 2025-11-29
 *
 * @copyright Copyright (c) 2025
 *
 */
#define BOOST_TEST_MODULE TestKernel
#include <CL/cl.h>

#include <boost/filesystem.hpp>
#include <boost/format.hpp>
#include <boost/format/format_fwd.hpp>
#include <boost/smart_ptr/scoped_array.hpp>
#include <boost/test/included/unit_test.hpp>
#include <boost/test/tools/old/interface.hpp>
#include <boost/test/unit_test_log.hpp>
#include <boost/timer/timer.hpp>
#include <cstddef>
#include <string>
#include <vector>

#include "test_fixture.hpp"
#include "utility.hpp"

using boost::format;
// !scoped_ptr用于关于单个对象，使用new分配；scoped_array用于关于数组，使用new[]分配
using boost::scoped_array;

constexpr size_t kNumElements = 1024;

class KernelTestFixture : public CLTestFixture {
protected:
    KernelTestFixture() : CLTestFixture()
    {
        buffer_a_ = clCreateBuffer(context_,
                                   CL_MEM_READ_WRITE,
                                   kNumElements * sizeof(int),
                                   NULL,
                                   NULL);
        BOOST_TEST_REQUIRE(buffer_a_);
        buffer_b_ = clCreateBuffer(context_,
                                   CL_MEM_READ_WRITE,
                                   kNumElements * sizeof(int),
                                   NULL,
                                   NULL);
        BOOST_TEST_REQUIRE(buffer_b_);
        buffer_c_ = clCreateBuffer(context_,
                                   CL_MEM_READ_WRITE,
                                   kNumElements * sizeof(int),
                                   NULL,
                                   NULL);
    }

    ~KernelTestFixture() = default;

    cl_mem buffer_a_;
    cl_mem buffer_b_;
    cl_mem buffer_c_;
};

BOOST_FIXTURE_TEST_CASE(VectorAdd, KernelTestFixture)
{
    cl_int err = CL_SUCCESS;

    std::string source = LoadProgramFile("vector_add.cl");

    cl_program program = cl_.CreateProgram(source.c_str(), source.size());
    BOOST_CHECK(program);
    cl_kernel kernel = cl_.CreateKernel("vector_add");
    BOOST_CHECK(kernel);

    scoped_array<int> a(new int[kNumElements]);
    scoped_array<int> b(new int[kNumElements]);
    scoped_array<int> c(new int[kNumElements]);

    GeneratorRandom(a.get(), kNumElements);
    GeneratorRandom(b.get(), kNumElements);

    cl_event wait_events[2]{NULL};

    err = clEnqueueWriteBuffer(queue_,
                               buffer_a_,
                               CL_FALSE,
                               0,
                               kNumElements * sizeof(int),
                               a.get(),
                               0,
                               NULL,
                               &wait_events[0]);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    err = clEnqueueWriteBuffer(queue_,
                               buffer_b_,
                               CL_FALSE,
                               0,
                               kNumElements * sizeof(int),
                               b.get(),
                               0,
                               NULL,
                               &wait_events[1]);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &buffer_a_);
    BOOST_CHECK(err == CL_SUCCESS);
    err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &buffer_b_);
    BOOST_CHECK(err == CL_SUCCESS);
    err = clSetKernelArg(kernel, 2, sizeof(cl_mem), &buffer_c_);
    BOOST_CHECK(err == CL_SUCCESS);

    boost::timer::cpu_timer timer;

    size_t global_work_size = kNumElements;

    err = clEnqueueNDRangeKernel(
        queue_, kernel, 1, NULL, &global_work_size, NULL, 2, wait_events, NULL);
    BOOST_CHECK(err == CL_SUCCESS);

    clFinish(queue_);
    // std::cout << "kernel time: " << timer.format();
    BOOST_TEST_MESSAGE("kernel time: " << timer.format());

    clEnqueueReadBuffer(queue_,
                        buffer_c_,
                        CL_TRUE,
                        0,
                        kNumElements * sizeof(int),
                        c.get(),
                        0,
                        NULL,
                        NULL);

    clFinish(queue_);

    for (size_t i = 0; i < kNumElements; ++i)
    {
        BOOST_CHECK_EQUAL(c[i], a[i] + b[i]);
    }
}

BOOST_FIXTURE_TEST_CASE(Random, KernelTestFixture)
{
    cl_int err = CL_SUCCESS;

    std::string source  = LoadProgramFile("random.cl");
    cl_program  program = cl_.CreateProgram(source.c_str(), source.size());
    BOOST_REQUIRE(program);

    cl_kernel kernel = cl_.CreateKernel("fill_random");
    BOOST_REQUIRE(kernel);

    scoped_array<int> a(new int[kNumElements]);
    bzero(a.get(), kNumElements * sizeof(int));

    err = clEnqueueWriteBuffer(queue_,
                               buffer_a_,
                               CL_TRUE,
                               0,
                               kNumElements * sizeof(int),
                               a.get(),
                               0,
                               NULL,
                               NULL);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    // arg1: buffer
    err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &buffer_a_);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    // arg2: seed
    unsigned int seed = time(NULL);
    err               = clSetKernelArg(kernel, 1, sizeof(cl_int), &seed);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    // arg3: start
    int start = 1024;
    err       = clSetKernelArg(kernel, 2, sizeof(cl_int), &start);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    // arg4: end
    int end = 2048;
    err     = clSetKernelArg(kernel, 3, sizeof(cl_int), &end);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    {
        SCOPED_TIMER("Fill Random Kernel");
        size_t global_work_size = kNumElements;
        err                     = clEnqueueNDRangeKernel(
            queue_, kernel, 1, NULL, &global_work_size, NULL, 0, NULL, NULL);
        BOOST_CHECK_EQUAL(err, CL_SUCCESS);

        clFinish(queue_);
    }

    {
        SCOPED_TIMER("Read Buffer");
        clEnqueueReadBuffer(queue_,
                            buffer_a_,
                            CL_TRUE,
                            0,
                            kNumElements * sizeof(int),
                            a.get(),
                            0,
                            NULL,
                            NULL);
    }

    for (size_t i = 0; i < kNumElements; ++i)
    {
        BOOST_CHECK((a[i] >= start) && (a[i] <= end));
    }
}

// 测试使用clCreateKernel创建内核
BOOST_AUTO_TEST_CASE(TestCreateKernel)
{
    CLWrapper cl;
    bool      init_result = cl.Init();
    BOOST_CHECK(init_result == true);

    cl_context   context = cl.GetContext();
    cl_device_id device  = cl.GetDeviceID();
    BOOST_CHECK(context != nullptr);
    BOOST_CHECK(device != nullptr);

    // 简单的内核源码用于测试
    const char* kernel_source =
        "__kernel void test_kernel(__global int* data, const int multiplier) { "
        "    int idx = get_global_id(0); "
        "    data[idx] *= multiplier; "
        "} "
        "__kernel void second_kernel(__global int* input, __global int* "
        "output) { "
        "    int idx = get_global_id(0); "
        "    output[idx] = input[idx] * input[idx]; "
        "}";

    // 创建程序
    cl_int     err = CL_SUCCESS;
    cl_program program =
        clCreateProgramWithSource(context, 1, &kernel_source, nullptr, &err);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK(program != nullptr);

    // 构建程序
    err = clBuildProgram(program, 1, &device, nullptr, nullptr, nullptr);
    if (err != CL_SUCCESS)
    {
        size_t log_size = 0;
        clGetProgramBuildInfo(
            program, device, CL_PROGRAM_BUILD_LOG, 0, nullptr, &log_size);
        if (log_size > 1)
        {
            std::vector<char> build_log(log_size);
            clGetProgramBuildInfo(program,
                                  device,
                                  CL_PROGRAM_BUILD_LOG,
                                  log_size,
                                  build_log.data(),
                                  nullptr);
            std::cout << "Build log: " << build_log.data() << std::endl;
        }
    }
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    // 创建内核
    cl_kernel kernel = clCreateKernel(program, "test_kernel", &err);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK(kernel != nullptr);

    // 验证内核
    cl_program kernel_program = nullptr;
    err                       = clGetKernelInfo(kernel,
                          CL_KERNEL_PROGRAM,
                          sizeof(kernel_program),
                          &kernel_program,
                          nullptr);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK(kernel_program == program);

    cl_context kernel_context = nullptr;
    err                       = clGetKernelInfo(kernel,
                          CL_KERNEL_CONTEXT,
                          sizeof(kernel_context),
                          &kernel_context,
                          nullptr);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK(kernel_context == context);

    // 检查内核名称
    char   kernel_name[256];
    size_t name_size = 0;
    err              = clGetKernelInfo(kernel,
                          CL_KERNEL_FUNCTION_NAME,
                          sizeof(kernel_name),
                          kernel_name,
                          &name_size);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK_EQUAL(std::string(kernel_name), "test_kernel");
    BOOST_CHECK_GT(name_size, 0);

    // 清理
    err = clReleaseKernel(kernel);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    err = clReleaseProgram(program);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    cl.DeInit();
}

// 测试内核工作组信息查询
BOOST_AUTO_TEST_CASE(TestKernelWorkGroupInfo)
{
    CLWrapper cl;
    bool      init_result = cl.Init();
    BOOST_CHECK(init_result == true);

    cl_context   context = cl.GetContext();
    cl_device_id device  = cl.GetDeviceID();
    BOOST_CHECK(context != nullptr);
    BOOST_CHECK(device != nullptr);

    // 简单的向量加法内核代码
    const char* vector_add_kernel =
        "__kernel void vector_add(__global const int* a, "
        "                          __global const int* b, "
        "                          __global int* c) { "
        "    int idx = get_global_id(0); "
        "    c[idx] = a[idx] + b[idx]; "
        "}";

    // 创建程序
    cl_int     err     = CL_SUCCESS;
    cl_program program = clCreateProgramWithSource(
        context, 1, &vector_add_kernel, nullptr, &err);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK(program != nullptr);

    // 构建程序
    err = clBuildProgram(program, 1, &device, nullptr, nullptr, nullptr);
    if (err != CL_SUCCESS)
    {
        size_t log_size = 0;
        clGetProgramBuildInfo(
            program, device, CL_PROGRAM_BUILD_LOG, 0, nullptr, &log_size);
        if (log_size > 1)
        {
            std::vector<char> build_log(log_size);
            clGetProgramBuildInfo(program,
                                  device,
                                  CL_PROGRAM_BUILD_LOG,
                                  log_size,
                                  build_log.data(),
                                  nullptr);
            std::cout << "Build log: " << build_log.data() << std::endl;
        }
    }
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    // 创建内核
    cl_kernel kernel = clCreateKernel(program, "vector_add", &err);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK(kernel != nullptr);

    // 查询内核工作组信息
    size_t work_group_size = 0;
    err                    = clGetKernelWorkGroupInfo(kernel,
                                   device,
                                   CL_KERNEL_WORK_GROUP_SIZE,
                                   sizeof(work_group_size),
                                   &work_group_size,
                                   nullptr);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    BOOST_CHECK_GT(work_group_size, 0);

    size_t compile_work_group_size[3];
    err = clGetKernelWorkGroupInfo(kernel,
                                   device,
                                   CL_KERNEL_COMPILE_WORK_GROUP_SIZE,
                                   sizeof(compile_work_group_size),
                                   &compile_work_group_size,
                                   nullptr);
    // 这个查询可能返回0值，这在某些内核中是正常的

    size_t preferred_work_group_size_multiple = 0;
    err                                       = clGetKernelWorkGroupInfo(kernel,
                                   device,
                                   CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE,
                                   sizeof(preferred_work_group_size_multiple),
                                   &preferred_work_group_size_multiple,
                                   nullptr);
    if (err == CL_SUCCESS)
    {
        BOOST_CHECK_GT(preferred_work_group_size_multiple, 0);
    }
    // 注意：这个查询在某些实现中可能不支持

    cl_ulong local_mem_size = 0;
    err                     = clGetKernelWorkGroupInfo(kernel,
                                   device,
                                   CL_KERNEL_LOCAL_MEM_SIZE,
                                   sizeof(local_mem_size),
                                   &local_mem_size,
                                   nullptr);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    // 本地内存大小可能为0

    cl_ulong private_mem_size = 0;
    err                       = clGetKernelWorkGroupInfo(kernel,
                                   device,
                                   CL_KERNEL_PRIVATE_MEM_SIZE,
                                   sizeof(private_mem_size),
                                   &private_mem_size,
                                   nullptr);
    if (err == CL_SUCCESS)
    {
        // 专用内存大小可能为0或非0
    }

    // 清理
    err = clReleaseKernel(kernel);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);
    err = clReleaseProgram(program);
    BOOST_CHECK_EQUAL(err, CL_SUCCESS);

    cl.DeInit();
}